﻿
/**
------------------------------------------------------------
* NGUYEN SY BACH
* bit.mta@gmail.com
* Hoc Vien Ky Thuat Quan Su
------------------------------------------------------------
*** Bien doi su dung cuda
------------------------------------------------------------
*/

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <assert.h>
#include "nsb_helper.h"
#include "device_launch_parameters.h"


__device__ inline void NearestNeighbor_dev(const uchar *inPtr, int strideIn,
	int channels, double x, double y,
	uchar *outPtr, int strideOut, int u, int v)
{
	x = int(x + 0.5);
	y = int(y + 0.5);
	int i = INDEX(int(x), int(y), strideIn, channels);
	int j = INDEX(u, v, strideOut, channels);
	for (int r = 0; r < channels; ++r){
		outPtr[j + r] = inPtr[i + r];
	}
}


__global__ void TransformKernel(
	uchar *outPtr, int u0, int v0,
	int outHeight, int outWidth, int channels,
	const double *g,
	const uchar *inPtr, int x0, int y0,
	int inHeight, int inWidth,
	int mode)
{
	double x;
	double y;
	int u;
	int v;
	v = threadIdx.y + blockIdx.y * blockDim.y;
	while (v < outHeight){
		u = threadIdx.x + blockIdx.x * blockDim.x;
		while (u < outWidth){
			x = g[0] * (u0 + u) + g[1] * (v0 + v) + g[2];
			y = g[3] * (u0 + u) + g[4] * (v0 + v) + g[5];
			x -= x0, y -= y0;
			if (x >= 0 && x < inWidth && y >= 0 && y < inHeight){
				NearestNeighbor_dev(inPtr, channels*inWidth, channels, x, y,
					outPtr, channels*inHeight, u, v);
			}
			else{
				int i = INDEX(u, v, channels*outWidth, channels);
				outPtr[i] = outPtr[i + 1] = outPtr[i + 2] = 0;
			}
			u += gridDim.x * blockDim.x;
		}
		v += gridDim.y * blockDim.y;
	}
}



#ifdef __cplusplus 
extern "C"
{
#endif

	//---------- Hàm biến đổi sử dụng GPU ----------
	// outPtr: con trỏ của mẩu ảnh kết quả
	// u0: Tọa độ theo cột của mẩu ảnh kết quả
	// v0: Tạo độ theo hàng của mẩu ảnh kết quả
	// outWidth: chiều rộng của mẩu ảnh kết quả
	// outHeight: chiều cao của mẩu ảnh kết quả
	// g: tham số của hàm ánh xạ ngược
	// inPtr: con trỏ của mẩu ảnh gốc
	// x0: Tạo độ theo cột của mẩu ảnh gốc
	// y0: Tạo độ theo hàng của mẩu ảnh gốc
	// inWidth: chiều rộng của mẩu ảnh gốc
	// inHeight: chiều cao của mẩu ảnh gốc
	// mode: Phương pháp tái chế mẫu
	// threadsPerBlock: Số theads 2D trên một Blocks, mặc định (4,4)
	// blocksPerGrid: Số blocks 2D trên Grid, mặc định (4,4)
	void TransformWithGPU(BYTE *&outPtr, int u0, int v0,
		int outWidth, int outHeight, int channels, const double *g,
		const BYTE *inPtr, int x0, int y0,
		int inWidth, int inHeight, int mode,
		dim3 threadsPerBlock = dim3(32,32),
		dim3 blocksPerGrid = dim3(16,16))
	{
		HANDLE_NULL(inPtr);
		HANDLE_NULL(outPtr);
		uchar *inPtrDev = NULL; // con trỏ mẩu ảnh gốc trên device
		uchar *outPtrDev = NULL; // con trỏ mẩu ảnh kết quả trên device
		double *gPtrDev = NULL; // tham số của hàm ánh xạ ngược trên device
		// Cấp phát bộ nhớ
		int mem_size = inHeight * inWidth * channels *sizeof(BYTE);
		int mem_size_out = outHeight * outWidth * channels *sizeof(BYTE);

		HANDLE_ERROR(cudaMalloc((void**)&inPtrDev,mem_size));
		HANDLE_ERROR(cudaMalloc((void**)&outPtrDev,mem_size_out));
		HANDLE_ERROR(cudaMalloc((void**)&gPtrDev,
			6 * sizeof(double)));
		//Copy dữ liệu từ host sang device
		printf("\nmem_size=%d mem_size_out=%d\n"
			"inHeight=%d inWidth=%d\n"
			"outHeight=%d outWidth=%d\n",
			mem_size, mem_size_out,
			inHeight, inWidth, 
			outHeight, outWidth);

		HANDLE_ERROR(cudaMemcpy(inPtrDev, inPtr,
			inWidth*inHeight*channels,
			cudaMemcpyHostToDevice));
		HANDLE_ERROR(cudaMemcpy(gPtrDev, g,
			6 * sizeof(double),
			cudaMemcpyHostToDevice));
		// chạy nhân
		TransformKernel << < blocksPerGrid, threadsPerBlock >> > (
			outPtrDev, u0, v0, outHeight, outWidth,
			channels, gPtrDev,
			inPtrDev, x0, y0, inHeight, inWidth,
			mode);

		HANDLE_ERROR(cudaGetLastError());
		HANDLE_ERROR(cudaMemcpy(outPtr, outPtrDev,
			mem_size_out, cudaMemcpyDeviceToHost));

		cudaFree(inPtrDev);
		cudaFree(outPtrDev);
		cudaFree(gPtrDev);
	}



#ifdef __cplusplus
}
#endif
